#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2;
template<typename T> class KernelLogSumExp {
public:
    __aicore__ inline KernelLogSumExp() {}
    __aicore__ inline void Init(GM_ADDR x,  GM_ADDR y, 
                            uint32_t blockLength,
                              uint32_t tileNum, uint32_t tileLength,
                              uint32_t lasttileLength, uint32_t typeKey,
                               uint32_t totalLength,
                              uint32_t stepSize, uint32_t unitCount) {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        this->typeKey = typeKey;
        if (this->typeKey == 0)
        {
            this->yLength = 8;
        }
        else
        {
            this->yLength = 16;
        }
        this->totalLength = totalLength;
        this->stepSize = stepSize;
        this->unitCount = unitCount;
        this->blockLength = blockLength;
        this->tileNum =
            tileNum ASSERT(tileNum != 0 && "tile num can not be zero!");
        this->tileLength = tileLength;
        this->lasttileLength = lasttileLength;
        uint32_t xLen = this->totalLength * this->stepSize;
        xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x,
                            blockLength > xLen ? blockLength : xLen);
        yGm.SetGlobalBuffer((__gm__ DTYPE_X *)y,
                            blockLength > xLen ? blockLength : xLen);
        yGm_f.SetGlobalBuffer((__gm__ float *)y,
                            blockLength > xLen ? blockLength : xLen);
      //   if (stepSize == 1 && typeKey == 0 ){
      //       this->bufferNum = 2;
      // this->tileLength = tileLength / this->bufferNum;
      // if (this->tileNum * 2 % 8 == 0)
      // {
      //   this->tileNumAlign = this->tileNum * 2;
      // }
      // else
      // {
      //   this->tileNumAlign = (this->tileNum * 2 + 7) / 8 * 8;
      // }
      // if (this->lasttileLength % 16 == 0)
      // {
      //   this->lasttileLengthAlign = this->lasttileLength;
      // }
      // else
      // {
      //   this->lasttileLengthAlign = (this->lasttileLength + 15) / 16 * 16;
      // }
      // pipe.InitBuffer(inQueueX_DB, this->bufferNum, this->tileLength * sizeof(float));
      // pipe.InitBuffer(inQueueT_DB, this->bufferNum, this->tileLength * sizeof(float));
      // pipe.InitBuffer(outQueueY_DB, this->bufferNum, this->yLength * sizeof(float));
      // pipe.InitBuffer(calcBuf, this->tileLength * sizeof(float));
      // pipe.InitBuffer(calcBuf3, this->tileLength * sizeof(float));
      //pipe.InitBuffer(calcBuf3,this->tileLength * sizeof(float));
    // }
    // else
    // {
      this->bufferNum = 1;
      pipe.InitBuffer(inQueueX, this->bufferNum, this->tileLength * sizeof(DTYPE_X));
      pipe.InitBuffer(inQueueT, this->bufferNum, this->yLength * sizeof(float));
      pipe.InitBuffer(outQueueY, this->bufferNum, this->yLength * sizeof(DTYPE_X));
      pipe.InitBuffer(outQueueY_F, this->bufferNum, this->yLength * sizeof(float));
      pipe.InitBuffer(calcBuf, this->bufferNum * this->tileLength * sizeof(float));

      if constexpr(std::is_same_v<T, half>){

        pipe.InitBuffer(xfBuf, this->bufferNum * this->tileLength * sizeof(float));

      }
      
      //pipe.InitBuffer(calcBuf3, this->bufferNum * this->tileLength * sizeof(float));
     // pipe.InitBuffer(calcBuf2, this->bufferNum * this->tileLength * sizeof(float));
    //   pipe.InitBuffer(calcBuf1, this->bufferNum * this->tileLength * sizeof(uint8_t));
    //   pipe.InitBuffer(calcBuf2, this->bufferNum * this->tileLength * sizeof(float));
    //}
        }   //// 此类特殊情况，可以适用double buffer处理进行优化

    
    
    __aicore__ inline void Process() {

        if (stepSize != 1)
    {
      // 处理多轴情况
        ProcessAxes();
        return;
    }

    // 处理一般情况
    int32_t loopCount = this->tileNum * this->bufferNum;
    for (int32_t i = 0; i < loopCount; i++)
    {
      CopyIn(i);
      Compute(i);
      CopyOut(i);
    }

    if (this->typeKey == 1)
      {
        // ========CopyIn======
        LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
        DataCopy(tLocal, yGm_f[0], this->yLength);
        inQueueT.EnQue(tLocal);
        // ========Compute======
        tLocal = inQueueT.DeQue<float>();
        LocalTensor<DTYPE_X> yLocal = outQueueY.AllocTensor<DTYPE_X>();
        Cast(yLocal, tLocal, AscendC::RoundMode::CAST_NONE, this->yLength);
        inQueueT.FreeTensor(tLocal);
        outQueueY.EnQue<DTYPE_X>(yLocal);
        // ========CopyOut======
        yLocal = outQueueY.DeQue<DTYPE_X>();
        DataCopy(yGm[0], yLocal, this->yLength);
        outQueueY.FreeTensor(yLocal);
      }
    else
      {
        // ========CopyIn======
        LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
        DataCopy(tLocal, yGm_f[0], this->yLength);
        inQueueT.EnQue(tLocal);
        // ========Compute======
        tLocal = inQueueT.DeQue<float>();
        LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
        Adds(yLocal, tLocal, (float)0.0, this->yLength);
        inQueueT.FreeTensor(tLocal);
        outQueueY.EnQue<float>(yLocal);
        // ========CopyOut======
        yLocal = outQueueY.DeQue<float>();
        DataCopy(yGm_f[0], yLocal, this->yLength);
        outQueueY.FreeTensor(yLocal);
      }
    }

        
    
private:
   __aicore__ inline void ProcessAxes()
  {
    int32_t loopCount = this->tileNum;
    for (int32_t k = 0; k < this->unitCount; k++)
    {
      for (int32_t j = 0; j < this->stepSize; j++)
      {
        for (int32_t i = 0; i < loopCount; i++)
        {
          CopyInAxes(i, j, k);
          Compute(i);
          CopyOutAxes(i, j, k);
        }
        if (this->typeKey == 1)
          {
            // ========CopyIn======
            LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
            DataCopy(tLocal, yGm_f[j + k * this->stepSize], this->yLength);
            inQueueT.EnQue(tLocal);
            // ========Compute======
            tLocal = inQueueT.DeQue<float>();
            LocalTensor<DTYPE_X> yLocal = outQueueY.AllocTensor<DTYPE_X>();
            Cast(yLocal, tLocal, AscendC::RoundMode::CAST_NONE, this->yLength);
            inQueueT.FreeTensor(tLocal);
            outQueueY.EnQue<DTYPE_X>(yLocal);
            // ========CopyOut======
            yLocal = outQueueY.DeQue<DTYPE_X>();
            DataCopy(yGm[j + k * this->stepSize], yLocal, this->yLength);
            outQueueY.FreeTensor(yLocal);
          }
        else
          {
            // ========CopyIn======
            LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
            DataCopy(tLocal, yGm_f[j + k * this->stepSize], this->yLength);
            inQueueT.EnQue(tLocal);
            // ========Compute======
            tLocal = inQueueT.DeQue<float>();
            LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
            Adds(yLocal, tLocal, (float)0.0, this->yLength);
            inQueueT.FreeTensor(tLocal);
            outQueueY.EnQue<float>(yLocal);
            // ========CopyOut======
            yLocal = outQueueY.DeQue<float>();
            DataCopy(yGm_f[j + k * this->stepSize], yLocal, this->yLength);
            outQueueY.FreeTensor(yLocal);
          }
      }
    }
  }

   __aicore__ inline void CopyInAxes(int32_t progress, int j, int k)
  {
    LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
    LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
    uint32_t startIndex = j + k * this->stepSize * this->totalLength +
                          progress * tileLength * this->stepSize;
    if (progress == this->tileNum - 1)
    {
      for (int i = 0; i < lasttileLength; i++)
      {
        xLocal.SetValue(i, xGm.GetValue(startIndex + i * stepSize));
      }

      //DataCopy(xLocal,xGm[startIndex])
      for (int i = lasttileLength; i < this->tileLength; i++)
      {
         xLocal.SetValue(i, (DTYPE_X)0);
      //   // if (pType == 0 || pType == 3 || pType == 4 || pType == 5)
      //   // {
      //   //   xLocal.SetValue(i, (DTYPE_X)0);
      //   }
        // else
        // {
          
          //xLocal.SetValue(i, xLocal.GetValue(0));
        //}
      }
     }
    else
    {
    for (int i = 0; i < tileLength; i++)
      {
        xLocal.SetValue(i, xGm.GetValue(startIndex + i * stepSize));
      }
    }
    DataCopy(tLocal, yGm_f[j + k * this->stepSize], this->yLength);
    if (progress == 0)
    {
      tLocal.SetValue(0, (DTYPE_X)0);
    }
    inQueueT.EnQue(tLocal);
    inQueueX.EnQue(xLocal);
  }
  __aicore__ inline void CopyOutAxes(int32_t progress, int j, int k)
  {
    LocalTensor<float> yLocal = outQueueY_F.DeQue<float>();
    DataCopy(yGm_f[j + k * this->stepSize], yLocal, this->yLength);
    outQueueY_F.FreeTensor(yLocal);
  }
  __aicore__ inline void Compute(int32_t progress)
  {
    LocalTensor<float> yLocal = outQueueY_F.AllocTensor<float>(); //this->tileLength
    LocalTensor<float> tLocal = inQueueT.DeQue<float>(); // this->yLength
    LocalTensor<float> wLocal = calcBuf.Get<float>();  // this->tileLength
    //LocalTensor<float> cLocal = calcBuf3.Get<float>();  // this->tileLength
    //LocalTensor<float> xLocal_f = xfBuf.Get<float>();
    if (this->typeKey == 0)
    {
      LocalTensor<float> xLocal = inQueueX.DeQue<float>();
      //LocalTensor<float> xLocal_f = xLocal;
      ReduceMax(wLocal,xLocal,wLocal,this->tileLength,false);
      // //Duplicate(wLocal,wLocal.GetValue(0),this->tileLength);
      Muls(wLocal,wLocal,(float)(-1.0),1);

      Adds(xLocal,xLocal,wLocal.GetValue(0),this->tileLength);
      // //Sub(xLocal,xLocal,wLocal,this->tileLength);
      // Exp(xLocal,xLocal,this->tileLength);
      // //for(int32_t i=0; i<)
      // ReduceSum(xLocal,xLocal,cLocal,this->tileLength);
    
      Exp(xLocal,xLocal,this->tileLength);  
      ReduceSum(xLocal,xLocal,yLocal,this->tileLength);
      Ln(yLocal,xLocal,1);
      //Muls(wLocal,wLocal,(float)(-1.0),1);
      Sub(yLocal,yLocal,wLocal,1);
      Add(yLocal, yLocal, tLocal, this->yLength);





    //   if (pType == 0)
    //   {
        // sqrt(sum(abs(x)^2))
    //   }
      
      outQueueY_F.EnQue<float>(yLocal);
      inQueueX.FreeTensor(xLocal);
      inQueueT.FreeTensor(tLocal);
    }
    else
    {
      LocalTensor<half> xLocal = inQueueX.DeQue<half>();
      LocalTensor<float> xLocal_f = xfBuf.Get<float>();
      Cast(xLocal_f, xLocal, RoundMode::CAST_NONE, this->tileLength);

      ReduceMax(wLocal,xLocal_f,wLocal,this->tileLength,false);
      Duplicate(wLocal,wLocal.GetValue(0),this->tileLength);
      Sub(xLocal_f,xLocal_f,wLocal,this->tileLength);
      Exp(xLocal_f,xLocal_f,this->tileLength);
      //for(int32_t i=0; i<)
      ReduceSum(xLocal_f,xLocal_f,yLocal,this->tileLength);
      Ln(xLocal_f,xLocal_f,1);
      Adds(yLocal,xLocal_f,wLocal.GetValue(0),1);
      Add(yLocal, yLocal, tLocal, this->yLength);
    //   if (pType == 0)
      //{
        // sqrt(sum(abs(x)^2))
    //   AscendC::Abs(xLocal_f, xLocal_f, this->tileLength);
    //   AscendC::Mul(xLocal_f, xLocal_f, xLocal_f, this->tileLength);
    //   AscendC::ReduceSum(yLocal, xLocal_f, wLocal, this->tileLength);
    //   AscendC::Add(yLocal, yLocal, tLocal, this->yLength);
      //}
      
      outQueueY_F.EnQue<float>(yLocal);
      inQueueX.FreeTensor(xLocal);
      inQueueT.FreeTensor(tLocal);
    }
  }
    __aicore__ inline void CopyOut(int32_t progress)
  {
    LocalTensor<float> yLocal = outQueueY_F.DeQue<float>();
    DataCopy(yGm_f[0], yLocal, this->yLength);
    outQueueY_F.FreeTensor(yLocal);
  }
  __aicore__ inline void CopyIn(int32_t progress)
  {
    LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
    LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
    if (progress == this->tileNum - 1)
    {
      if (progress == 0)
      {
        DataCopy(xLocal[0], xGm[0], this->tileLength);
        for (int i = this->totalLength; i < this->tileLength; i++)
        {
        //   if (pType == 0 || pType == 3 || pType == 4 || pType == 5)
        //   {
            xLocal.SetValue(i, (DTYPE_X)0);
        //   }
        //   else
        //   {
        //     xLocal.SetValue(i, xLocal.GetValue(0));
        //   }
        }
      }
      else
      {
        DataCopy(
            xLocal[0],
            xGm[(progress - 1) * this->tileLength + this->lasttileLength],
            this->tileLength);
        for (int i = 0; i < this->tileLength - this->lasttileLength; i++)
        {
        //   if (pType == 0 || pType == 3 || pType == 4 || pType == 5)
        //   {
            xLocal.SetValue(i, (DTYPE_X)0);
        //   }
        //   else
        //   {
        //     xLocal.SetValue(i, xLocal.GetValue(this->tileLength - this->lasttileLength));
        //   }
        }
      }
    }
    else
    {
      DataCopy(xLocal[0], xGm[progress * this->tileLength],
               this->tileLength);
    }
    DataCopy(tLocal, yGm_f[0], this->yLength);
    if (progress == 0)
    {
      tLocal.SetValue(0, (DTYPE_X)0);
    }
    inQueueT.EnQue(tLocal);
    inQueueX.EnQue(xLocal);
  }
    



   
   
private:
  TPipe pipe;
  TQue<QuePosition::VECIN, 2> inQueueX, inQueueT;
  //TQue<QuePosition::VECIN, 2> inQueueX_DB, inQueueT_DB;
  TQue<QuePosition::VECOUT, 2> outQueueY, outQueueY_F;
  //TQue<QuePosition::VECOUT, 2> outQueueY_DB;
  TBuf<TPosition::VECCALC> calcBuf,xfBuf;
  GlobalTensor<DTYPE_X> xGm;
  GlobalTensor<float> yGm_f;
  GlobalTensor<DTYPE_Y> yGm;
  uint32_t blockLength;
  uint32_t tileNum;
  uint32_t tileNumAlign;
  uint32_t tileLength;
  uint32_t totalLength;
  uint32_t stepSize;
  uint32_t unitCount;
  uint32_t lasttileLength;
  uint32_t lasttileLengthAlign; // 64B对齐后的最后一个tile的长度
  uint32_t typeKey;
  uint32_t yLength; // 输出元素的最小搬运长度(32B)
  uint32_t bufferNum; // BUFFER_NUM, 为2时开启Double Buffer

};

template<typename TYPE_X, typename TYPE_Y> class KernelLogSumExpReduce2D {
    using T = TYPE_X;
public:
    __aicore__ inline KernelLogSumExpReduce2D() {}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, int64_t M, int64_t N, int64_t Z) {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        this->M = M; //行
        this->N = N; //列
        this->Z = Z; //循环次数
        // this->dim0 =dim0;
        this->row32 = (N+31)/32*32;
        //this->dim = 0;
        
        xGm.SetGlobalBuffer((__gm__ DTYPE_X*)x,  M * N * Z);
        
       
        yGm.SetGlobalBuffer((__gm__ DTYPE_X*)y,  M*Z);
            
        
        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->row32 * sizeof(DTYPE_X));
        pipe.InitBuffer(outQueueY, BUFFER_NUM,  N * sizeof(DTYPE_X));
        
        //pipe.InitBuffer(outQueueZ,BUFFER_NUM, N*sizeof(DTYPE_X) );
        // pipe.InitBuffer(inQueueX1,BUFFER_NUM, (M*N*Z) * sizeof(DTYPE_X));
        // pipe.InitBuffer(inQueueX2,BUFFER_NUM, (M*N*Z) * sizeof(DTYPE_X));
        //pipe.InitBuffer(SumQueue, this->row32 * sizeof(DTYPE_X));
        // if(this->dim0==0){
        //     pipe.InitBuffer(QueueBuff0, M * sizeof(float));
        //     pipe.InitBuffer(QueueBuff2, M * sizeof(float));
        
        //     pipe.InitBuffer(QueueBuff1, M * sizeof(float));

        // }
        // else{
        pipe.InitBuffer(QueueBuff0, N * sizeof(float));
        //pipe.InitBuffer(QueueBuff2, 1 * sizeof(float));
    
        pipe.InitBuffer(QueueBuff1, N * sizeof(float));
        // }
        if constexpr (std::is_same_v<DTYPE_X, half>) {
          pipe.InitBuffer(QueueBuff2, N * sizeof(float));
          pipe.InitBuffer(QueueBuff3, N * sizeof(half));
        }
        // pipe.InitBuffer(QueueBuff3, N * sizeof(DTYPE_X));
        // pipe.InitBuffer(QueueBuff4, N * sizeof(DTYPE_X));
        // if constexpr (std::is_same_v<DTYPE_X, half>){
            
        //     pipe.InitBuffer(QueueBuff0, N * sizeof(float));
            
        //     //pipe.InitBuffer(SumQueue, 1, this->row32 * sizeof(float));
        // //pipe.InitBuffer(QueueBuff3, this->row32* sizeof(float));
        // }
    }
    __aicore__ inline void Process() {
        //InitTensor();
        // if(this->dim0==0){
            
        //     TransposeShape();
            
        // }
        for(int32_t i = 0; i < Z; i ++)
        {
            
                 for(int32_t j = 0; j <M; j++)
            {
                CopyIn(i, j);
                Compute(i, j);
                
                
               
            }
           
        }
        //CopyOut( );
        
        
    }
    
private:
    
    //  __aicore__ inline void InitTensor()
    // {
          
    //       LocalTensor<DTYPE_X> zLocal = outQueueZ.AllocTensor<DTYPE_X>();
    //       DTYPE_X zero =  0;
    //       Duplicate(zLocal, zero, N);
    //       outQueueZ.EnQue<DTYPE_X>(zLocal);
            
    //     }
       
    
    __aicore__ inline void CopyIn(int32_t i, int32_t j)
    {
        LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
        // if(this->dim0==0){
        //     DataCopy(xLocal, xGm[i*M*N + j*M], this->row32);
        // }

        LocalTensor<DTYPE_X> yLocal = outQueueY.AllocTensor<DTYPE_X>();
        DTYPE_X zero =  0;
        Duplicate(yLocal, zero, N);
        outQueueY.EnQue<DTYPE_X>(yLocal);
        
        DataCopy(xLocal, xGm[i*M*N + j*N], this->row32);
        
        
        

        
        inQueueX.EnQue(xLocal);
    }
    __aicore__ inline void Compute(int32_t i, int32_t j)
    {
        LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
        
        LocalTensor<DTYPE_X> yLocal = outQueueY.DeQue<DTYPE_X>();
        //LocalTensor<DTYPE_X> zLocal = outQueueZ.DeQue<DTYPE_X>();
        
        
        
        if constexpr (std::is_same_v<DTYPE_X, half>){
            
         
            auto p1 = QueueBuff0.Get<float>();
            auto p2 = QueueBuff1.Get<float>();
            auto p3 = QueueBuff2.Get<float>();
            auto p4 = QueueBuff3.Get<DTYPE_X>();
            Cast(p1, xLocal, RoundMode::CAST_NONE, N);
           
            ReduceMax(p2,p1,p2,N,false);
            Duplicate(p3,p2.GetValue(0),N);
            Sub(p1,p1,p3,N);


            Exp(p1,p1,N);
            ReduceSum(p1,p1,p2,N);
            Ln(p1,p1,1);
            
            Adds(p1,p1,p3.GetValue(0),1);
            Cast(yLocal, p1, RoundMode::CAST_NONE, 1);
            //educeSum(zLocal[j],yLocal,p4,N);
            //outQueueY.EnQue<DTYPE_X>(yLocal);

            yGm.SetValue(j,yLocal.GetValue(0));
            outQueueY.FreeTensor(yLocal);

            //outQueueZ.EnQue(zLocal);
           
        }
        else{
            //LocalTensor<DTYPE_Y> zLocal = outQueueZ.DeQue<float>();

            //LocalTensor<DTYPE_X> tmpB2 = QueueBuff2.Get<DTYPE_X>();
            LocalTensor<DTYPE_X> tmpB3 = QueueBuff0.Get<DTYPE_X>();
            LocalTensor<DTYPE_X> tmpB4 = QueueBuff1.Get<DTYPE_X>();
           
            ReduceMax(tmpB3,xLocal,tmpB3,N,false);
            Duplicate(tmpB4,tmpB3.GetValue(0),N);
            Sub(xLocal,xLocal,tmpB4,N);


            Exp(xLocal,xLocal,N);
            ReduceSum(xLocal,xLocal,tmpB4,N);
            Ln(xLocal,xLocal,1);
            Adds(yLocal,xLocal,tmpB3.GetValue(0),1);


            yGm.SetValue(j,yLocal.GetValue(0));
            //ReduceSum(zLocal[j],yLocal,tmpB4,N);
           
            
            //Copy(zLocal,yLocal,1,1,{1,1,8,8});
            //` Maxs(yLocal,yLocal,static_cast<DTYPE_X>(0.0),1);
            //到这里是没问题的
            //outQueueY.EnQue<DTYPE_Y>(yLocal);
            //for(int32_t i=0; i<N;i++)
            
            //DataCopy(yGm[j])
            //outQueueZ.FreeTensor(zLocal);

            outQueueY.FreeTensor(yLocal);


            //outQueueY.EnQue<DTYPE_X>(yLocal);

            //outQueueZ.EnQue(zLocal);
            

        }
        
        //inQueueX.EnQue(xLocal);
        inQueueX.FreeTensor(xLocal);
       
        
        
    }
    // __aicore__ inline void CopyOut()
    // {   
    //     // if constexpr (std::is_same_v<DTYPE_X, half>){
        
    //    // LocalTensor<DTYPE_Y> yLocal = outQueueY.DeQue<DTYPE_X>();
    //     // LocalTensor<DTYPE_X> tmpB1 = QueueBuff0.Get<DTYPE_X>();
    //     // //LocalTensor<DTYPE_X> tmpB2 = QueueBuff2.Get<DTYPE_X>();
    //     // LocalTensor<DTYPE_X> tmpB4 = QueueBuff1.Get<DTYPE_X>();
    //     //
        
    //     LocalTensor<DTYPE_X> zLocal = outQueueZ.DeQue<DTYPE_X>();

        
    //     //Copy(yLocal,zLocal,N);

    //     // for(int i=0;i<N;i++){

    //     //     yLocal.SetValue(i,zLocal.GetValue(i));
    //     // }
    //     DataCopy(yGm,zLocal,N);
          
    //     outQueueZ.FreeTensor(zLocal);
           
        
        
    //     //DataCopy(yGm,yLocal, 1);
    //     //outQueueY.FreeTensor(yLocal);
        
    //     //SumQueue.FreeTensor(sum);
    //     //inQueueX.FreeTensor(xLocal);
    // }
private:
    TPipe pipe;
    // TBuf<QuePosition::VECCALC> tmpBuffer, signbitBuffer;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX,inQueueX1;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
    //TQue<QuePosition::VECOUT, 1> SumQueue;
    //TBuf<QuePosition::VECCALC>  QueueBuff0,QueueBuff1,QueueBuff2;
    TBuf<QuePosition::VECCALC>  QueueBuff0,QueueBuff1,QueueBuff2,QueueBuff3;

    //TQue<QuePosition::VECCALC,BUFFER_NUM>  outQueueZ;
    
    GlobalTensor<DTYPE_X> xGm;
    GlobalTensor<DTYPE_Y> yGm;
    GlobalTensor<DTYPE_X> x1Gm;
    int64_t M;
    int64_t N;
    int64_t Z;
    int64_t dim0;
    bool keep_dim;
    int64_t row32;

    
};




extern "C" __global__ __aicore__ void log_sum_exp(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    // TODO: user kernel impl
    if(TILING_KEY_IS(1)){
      KernelLogSumExp<DTYPE_X> op;
      op.Init(x, y, tiling_data.blockLength,
            tiling_data.tileNum, tiling_data.tileLength,
            tiling_data.lasttileLength, tiling_data.typeKey,
          tiling_data.totalLength,
            tiling_data.stepSize, tiling_data.unitCount);
      op.Process();
    
    }
    else if(TILING_KEY_IS(2)){
      KernelLogSumExpReduce2D<DTYPE_X,DTYPE_Y> op;
      op.Init(x, y, tiling_data.M, tiling_data.N, tiling_data.Z);
      op.Process();


    }
   
}
